Fix parquet preprocess strings page bounds check for pages with zero rows - #23241

Related to NVIDIA/cudf-spark#15235.

Running the failing job with the fix.

_{Comment @coderabbitai help to get the list of available commands.}

Caution

Inline review comments failed to post. This is likely due to GitHub's internal server error or limits when posting large numbers of comments. If you are seeing this consistently it is likely a permissions issue. Please check "Moderation" -> "Code review limits" under your organization settings.

Actionable comments posted: 1

🤖 Prompt for all review comments with AI agents

Verify each finding against current code. Fix only still-valid issues, skip the
rest with a brief reason, keep changes minimal, and validate.

Inline comments:
In `@cpp/src/io/parquet/decode_preprocess.cu`:
- Around line 446-480: Initialize the shared copy_barrier exactly once before
the first staged-copy decoder pass, rather than calling init() again for the
definition pass. Update the repetition/definition flow around
decoders[level_type::REPETITION].init and decoders[level_type::DEFINITION].init
so both reuse the initialized barrier through arrive_and_wait(), preserving
synchronization between phases.

🪄 Autofix (Beta)
Fix all unresolved CodeRabbit comments on this PR:

• Push a commit to this branch (recommended)
• Create a new PR with the fixes

ℹ️ Review info
⚙️ Run configuration
Configuration used: Path: .coderabbit.yaml

Review profile: CHILL

Plan: Enterprise

Run ID: 0ab7872e-f737-4933-aae9-d3fdb1131958

📥 Commits
Reviewing files that changed from the base of the PR and between 2a6ca4e and 8148a66.

📒 Files selected for processing (4)

• cpp/src/io/parquet/decode_preprocess.cu
• cpp/src/io/parquet/page_string_decode.cu
• cpp/src/io/parquet/reader_impl_preprocess.cu
• cpp/tests/io/parquet_chunked_reader_test.cu

🚧 Files skipped from review as they are similar to previous changes (3)

• cpp/tests/io/parquet_chunked_reader_test.cu
• cpp/src/io/parquet/reader_impl_preprocess.cu
• cpp/src/io/parquet/page_string_decode.cu

🛑 Comments failed to post (1)

cpp/src/io/parquet/decode_preprocess.cu (1)> 446-480: 🩺 Stability & Availability | 🟠 Major | ⚡ Quick win

🧩 Analysis chain
🏁 Script executed:

#!/usr/bin/env bash
set -euo pipefail
rg -n -C 8 'cuda::barrier|copy_barrier|arrive_and_wait|init\(' cpp/src/io/parquet
rg -n -C 8 'class barrier|struct barrier|void init' . --glob '*barrier*'

Repository: rapidsai/cudf
Length of output: 28378

🏁 Script executed:

#!/usr/bin/env bash
set -euo pipefail

sed -n '200,260p' cpp/src/io/parquet/rle_stream.cuh
printf '\n----\n'
rg -n -C 3 'init\(&copy_barrier|arrive_and_wait\(\)' cpp/src/io/parquet
printf '\n----\n'
rg -n -C 3 'cuda::barrier<cuda::thread_scope_block>|init\(&.*barrier|arrive_and_wait\(\)' cpp/src -g '*.{cu,cuh,cpp,h,hpp}'

Repository: rapidsai/cudf
Length of output: 9988

🌐 Web query:
CUDA cuda::barrier reinitialize init arrive_and_wait semantics shared barrier phase reuse
💡 Result:
In the CUDA programming model, the cuda::barrier object is a reusable synchronization mechanism that manages thread coordination across a sequence of phases [1][2][3]. Initialization and Reinitialization Before any thread participates in a barrier, it must be explicitly initialized using the cuda::barrier::init() friend function [4][3]. For shared barriers, this initialization must happen before any participating thread arrives at the barrier [3][5]. Because initialization itself is a synchronization-dependent task, threads typically use a collective operation like block.sync() or syncthreads() to bootstrap the barrier before its first use [3][5]. There is no distinct reinitialize function for cuda::barrier. Instead, the barrier is designed to be self-resetting [3]. When the expected arrival count (the number of arrive() calls) reaches zero, the barrier automatically and atomically resets its internal counter to the expected arrival count (adjusted by any arrive_and_drop calls) and increments the phase [2][3]. Semantics and Phase Reuse The cuda::barrier lifetime consists of a sequence of phases [2]. This phase-based design is what enables reusability [2][6]. - arrive: This function decrements the barrier's expected arrival count [2][6]. It returns an arrival_token, which represents the current phase and must be passed to wait() [2][7]. - arrive_and_wait: This is a convenience function that performs the arrival and immediately waits on the same phase [8]. - Phase Completion: Once all expected threads have called arrive (or arrive_and_drop), the barrier completes the current phase, executes any associated completion function, and transitions to the next phase [2][3][6]. - Thread Wait: A thread that calls arrive can defer waiting on the phase synchronization point [2][6]. This allows threads to perform independent work between signaling their arrival and blocking for the phase to complete, which is a key mechanism for hiding synchronization latency [2][9]. Crucially, the barrier is only ready to be reused for the next phase once the current phase's completion step has finished [3]. If a thread's arrive() call triggers the countdown to zero, it must ensure it does not attempt to reuse the barrier for another arrive() call until the phase has transitioned [3]. In a correctly synchronized program, the automatic phase increment ensures that subsequent arrive() calls are correctly associated with the next phase [2][3].
Citations:

• 1: https://nvidia.github.io/cccl/unstable/libcudacxx/extended_api/synchronization_primitives/barrier.html
• 2: https://gevtushenko.github.io/cccl/libcudacxx/extended_api/synchronization_primitives/barrier.html
• 3: https://docs.nvidia.com/cuda/archive/13.2.1/cuda-programming-guide/04-special-topics/async-barriers.html
• 4: https://docs.nvidia.com/cuda/cuda-programming-guide/04-special-topics/async-barriers.html
• 5: https://docs.nvidia.com/cuda/archive/13.2.0/cuda-programming-guide/04-special-topics/async-barriers.html
• 6: https://cppreference.dev/w/cpp/thread/barrier
• 7: https://docs.nvidia.com/cuda/cuda-programming-guide/04-special-topics/cooperative-groups.html
• 8: https://github.com/NVIDIA/cccl/blob/main/libcudacxx/include/cuda/std/__barrier/barrier.h
• 9: https://docs.nvidia.com/cuda/cuda-programming-guide/05-appendices/device-callable-apis.html

🌐 Web query:
NVIDIA libcudacxx cuda::barrier init shared barrier may be called again same object after phase completion
💡 Result:
In the libcudacxx cuda::barrier API, you should not call the init function multiple times on the same object after phase completion [1][2][3]. According to official NVIDIA documentation, the barrier must be initialized using the cuda::barrier::init friend function before any thread participates [1][2]. Once the barrier is in use, it is designed to manage phases automatically [2]. When the last call to arrive occurs (completing the expected count), the barrier automatically and atomically resets its countdown for the next phase [2]. Re-initializing a barrier that is already in use or has completed a phase can lead to undefined behavior, as the internal state—including the phase bit and arrival count—is managed by the arrival and wait operations to ensure continuous, multi-phase synchronization [4][2]. If you need to reset the barrier or change its expected arrival count, you should not re-initialize the object itself; instead, you should manage the barrier lifecycle according to the specific scope requirements (such as thread_scope_block or thread_scope_device) [2][5]. For example, if a barrier is intended to be used across different, non-concurrent thread groups, ensure that any necessary memory synchronization or reuse logic is handled at the application level without attempting to re-run the init constructor on the same memory location while other threads might still be observing it [5].
Citations:

• 1: https://docs.nvidia.com/cuda/cuda-programming-guide/04-special-topics/async-barriers.html
• 2: https://docs.nvidia.com/cuda/archive/13.2.1/cuda-programming-guide/04-special-topics/async-barriers.html
• 3: https://docs.nvidia.com/cuda/archive/13.1.0/cuda-programming-guide/04-special-topics/async-barriers.html
• 4: https://github.com/NVIDIA/cccl/blob/main/libcudacxx/include/cuda/std/__barrier/barrier.h
• 5: NVIDIA/libcudacxx#325

🌐 Web query:
CUDA barrier init function semantics reused barrier phases docs
💡 Result:
In CUDA, a cuda::barrier object requires explicit initialization before any thread can participate in it [1][2]. The initialization must occur using the cuda::barrier::init() friend function, which defines the expected arrival count (the number of threads or operations that must arrive before the barrier completes its current phase) [1][2][3]. The lifecycle of a cuda::barrier consists of a sequence of barrier phases [4][5]. The semantics regarding initialization, phase advancement, and reuse are as follows: 1. Initialization: The init() function must be called before any thread arrives at the barrier [1][2]. Because this is a collective requirement, threads typically use a prior synchronization mechanism (such as syncthreads() or a cooperative group sync) to bootstrap the initialization [2]. 2. Phase Advancement: An asynchronous barrier maintains a countdown starting from the expected arrival count. As participating threads call arrive(), the countdown decrements [3][6]. When the countdown reaches zero, the barrier automatically and atomically resets the countdown to the original expected arrival count and advances to the next phase [3][6]. 3. Reuse and Sequencing: - Tokens: When a thread calls arrive(), it receives an arrival_token associated with the current phase [3][6]. - Blocking/Waiting: A call to wait(std::move(token)) blocks the thread if the barrier is still in the phase associated with that token [3][6]. If the phase has already advanced, the thread does not block [3][6]. - Reuse Constraint: If a thread's call to arrive() completes the countdown (reaching zero), the thread must call wait() before the barrier can be reused for a subsequent arrive() call [3][6]. - Token Validity: wait() must only be called using a token from the current phase or the immediately preceding phase; using tokens from other phases results in undefined behavior [3][6]. For cuda::barrier<cuda::thread_scope_block> objects residing in shared memory, a specific implementation constraint exists: they should only be used by the threads in the thread block (CTA) that constructed them [7]. Other barrier scopes or memory locations may allow for reuse across different thread groups, provided no concurrent accesses occur between the different groups [7].
Citations:

• 1: https://docs.nvidia.com/cuda/cuda-programming-guide/04-special-topics/async-barriers.html
• 2: https://docs.nvidia.com/cuda/archive/13.1.0/cuda-programming-guide/04-special-topics/async-barriers.html
• 3: https://docs.nvidia.com/cuda/archive/13.2.1/cuda-programming-guide/04-special-topics/async-barriers.html
• 4: https://nvidia.github.io/cccl/unstable/libcudacxx/extended_api/synchronization_primitives/barrier.html
• 5: https://gevtushenko.github.io/cccl/libcudacxx/extended_api/synchronization_primitives/barrier.html
• 6: https://docs.nvidia.com/cuda/archive/13.1.1/cuda-programming-guide/04-special-topics/async-barriers.html
• 7: NVIDIA/libcudacxx#325

Initialize copy_barrier once
cuda::barrier is reusable across phases; calling init() again on the same shared object is not the intended lifecycle and can trigger undefined behavior. Bootstrap it once before the first decoder that uses staged copies, then reuse the same barrier with arrive_and_wait() for the repetition and definition passes.

🤖 Prompt for AI Agents

Verify each finding against current code. Fix only still-valid issues, skip the
rest with a brief reason, keep changes minimal, and validate.

In `@cpp/src/io/parquet/decode_preprocess.cu` around lines 446 - 480, Initialize
the shared copy_barrier exactly once before the first staged-copy decoder pass,
rather than calling init() again for the definition pass. Update the
repetition/definition flow around decoders[level_type::REPETITION].init and
decoders[level_type::DEFINITION].init so both reuse the initialized barrier
through arrive_and_wait(), preserving synchronization between phases.

This has nothing to do with the current PR. This can be evaluated/handled elsewhere at an elsetime.

/merge